#include "CUDAFunctions.cuh"

texture<unsigned char,2,cudaReadModeElementType> texUniverse;
cudaChannelFormatDesc textureChannel = cudaCreateChannelDesc<unsigned char>();

//kernel function to compute the game of life progress
__global__ void tick(unsigned char* device_universe, unsigned char* device_next_universe,int rows,int columns)
{
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;

	if (ix > 0 && ix < rows - 1 && iy > 0 && iy < columns - 1)
	{
		int liveNeighbors = device_universe[(ix - 1) * columns + iy - 1] +
							device_universe[(ix - 1) * columns + iy] +
							device_universe[(ix - 1) * columns + iy + 1] +
							device_universe[ix * columns + iy - 1] +
							device_universe[ix * columns + iy + 1] +
							device_universe[(ix + 1) * columns + iy - 1] +
							device_universe[(ix + 1) * columns + iy] +
							device_universe[(ix + 1) * columns + iy + 1];
		
		if (device_universe[ix * columns + iy] == 1)
		{
			if (liveNeighbors < 2 || liveNeighbors > 3)
				device_next_universe[ix * columns + iy] = 0;
		}
		else
			if (liveNeighbors == 3)
				device_next_universe[ix * columns + iy] = 1;
	}
}

__global__ void update(unsigned char* device_universe,unsigned char* device_next_universe,int rows,int columns)
{
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;

	if (ix > 0 && ix < rows - 1 && iy > 0 && iy < columns - 1)
		device_universe[ix * columns + iy] = device_next_universe[ix * columns + iy];
}

__global__ void tickWithTexture(unsigned char* device_next_universe,int rows,int columns)
{
	int ix = blockIdx.x * blockDim.x + threadIdx.x;
	int iy = blockIdx.y * blockDim.y + threadIdx.y;

	//2D Texture has axis X and Y exchanged as opposed to standard cohordinates (i.e. X = vertical, Y = horizontal).
	//Therefore the universe matrix is transposed inside a 2D Texture
	if (ix > 0 && ix < rows - 1 && iy > 0 && iy < columns - 1)
	{
		int liveNeighbors =			    tex2D(texUniverse,iy - 1,ix - 1) +
										tex2D(texUniverse,iy - 1, ix) +
										tex2D(texUniverse,iy - 1,ix + 1) +
										tex2D(texUniverse,iy,ix - 1) +
										tex2D(texUniverse,iy,ix + 1) +
										tex2D(texUniverse,iy + 1,ix - 1) +
										tex2D(texUniverse,iy + 1,ix) +
										tex2D(texUniverse,iy + 1,ix + 1);

		if (tex2D(texUniverse,iy,ix) == 1)
		{
			if (liveNeighbors < 2 || liveNeighbors > 3)
				device_next_universe[ix * columns + iy] = 0;
		}
		else
			if (liveNeighbors == 3)
				device_next_universe[ix * columns + iy] = 1;
	}
}

void cudaCheckAlloc(cudaError_t alloc,char* message)
{
	if (alloc != cudaSuccess)
	{
		cout << message << endl;
		exit(1);
	}
}

void cudaRunWithTextureCache(bool save_to_file)
{
	rows = globalUniverse ->height();
	columns = globalUniverse ->width();
	unsigned char* host_universe = globalUniverse ->bufferAddress();

	size_t dev_size = rows * columns;
	dim3 threads(BLOCKDIM_X,BLOCKDIM_Y);
	dim3 grid((int)ceil((double)rows / (double)BLOCKDIM_X), (int)ceil((double)columns / (double)BLOCKDIM_Y));
	cudaError_t cuda_alloc;
	cudaArray* device_universe;
	unsigned char* device_next_universe;
	int devCount;
	cudaDeviceProp devProp;

	//system infos
	cout << "Scanning system for CUDA devices" << endl;
	cudaGetDeviceCount(&devCount);

	for (int i = 0; i < devCount; i++)
	{
		cudaGetDeviceProperties(&devProp,i);
		cout << "Found " << devCount << " devices" << endl;
		cout << "***** GPU SYSTEM INFORMATION: *****\n\n";
		cout << "Name: " << devProp.name << endl;
		cout << "Total global memory: " << devProp.totalGlobalMem << endl;
		cout << "Total block shared memory: " << devProp.sharedMemPerBlock << endl;
		cout << "Total registers per block: " << devProp.regsPerBlock << endl;
		cout << "Warp size: " << devProp.warpSize << endl;
		cout << "Maximum threads per block" << devProp.maxThreadsPerBlock << endl;

		for (int j = 0; j < 3; j++)
		{
			cout << "Maximum dimension " << j << " of block: " << devProp.maxThreadsDim[j] << endl;
			cout << "Maximum dimension " << j << " of grid: " << devProp.maxGridSize[j] << endl;
		}

		cout << "Clock rate: " << devProp.clockRate << endl;
		cout << "Total constant memory: " << devProp.totalConstMem << endl;
		cout << "Number of multiprocessors: " << devProp.multiProcessorCount << endl;
		cout << "\n" << endl;
	}

	//alocating memory on GPU
	cout << "Allocating memory on GPU..." << endl;
	cudaMallocArray(&device_universe,&textureChannel,rows,columns);
	cudaMalloc((void **)&device_next_universe,dev_size);
	cudaMemcpy(device_next_universe,host_universe,dev_size,cudaMemcpyHostToDevice);
	cudaMemcpyToArray(device_universe,0,0,device_next_universe,dev_size,cudaMemcpyDeviceToDevice);
	cudaBindTextureToArray(texUniverse,device_universe);
	texUniverse.filterMode = cudaFilterModePoint;
	texUniverse.normalized = 0;
	cout << "Done!" << endl;

	cout << "Running game of life on GPU..." << endl;
	for (int i = 0; i < TICK_COUNT; i++)
	{
		tickWithTexture <<< grid,threads >>>(device_next_universe,rows,columns);
		cudaDeviceSynchronize();
		cudaMemcpyToArray(device_universe,0,0,device_next_universe,rows * columns,cudaMemcpyDeviceToDevice);
	}
	cout << "Done!" << endl;
	//cudaMemcpyFromArray(host_universe,device_universe,0,0,dev_size,cudaMemcpyDeviceToHost);
	//cudaMemcpy(host_universe,device_next_universe,dev_size,cudaMemcpyDeviceToHost);

	cout << "Storing result in main memory..." << endl;
	cudaMemcpyFromArray(host_universe,device_universe,0,0,dev_size,cudaMemcpyDeviceToHost);
	cout << "Done" << endl;

	//release memory and unbind textures
	cudaUnbindTexture(texUniverse);
	cudaFreeArray(device_universe);
	cudaFree(device_next_universe);

	if (save_to_file)
	{
			cout << "Writing output file..." << endl;
			CreateDirectory(L"D:\\HPCTest",NULL);
			globalUniverse ->storeMatrix("D:\\HPCTest\\matrix.out");
			cout << "Done!" << endl;
	}

}

void cudaRunWithNormalMemory(bool save_to_file = false)
{
	unsigned char* host_universe = globalUniverse ->bufferAddress();
	rows = globalUniverse ->height();
	columns = globalUniverse ->width();
	size_t dev_size = rows * columns;
	dim3 threads(BLOCKDIM_X,BLOCKDIM_Y);
	dim3 grid((int)ceil((double)rows / (double)BLOCKDIM_X), (int)ceil((double)columns / (double)BLOCKDIM_Y));
	cudaError_t cuda_alloc;
	unsigned char* device_universe;
	unsigned char* device_next_universe;
	int devCount;
	cudaDeviceProp devProp;

	//system infos
	cout << "Scanning system for CUDA devices" << endl;
	cudaGetDeviceCount(&devCount);

	for (int i = 0; i < devCount; i++)
	{
		cudaGetDeviceProperties(&devProp,i);
		cout << "Found " << devCount << " devices" << endl;
		cout << "***** GPU SYSTEM INFORMATION: *****\n\n";
		cout << "Name: " << devProp.name << endl;
		cout << "Total global memory: " << devProp.totalGlobalMem << endl;
		cout << "Total block shared memory: " << devProp.sharedMemPerBlock << endl;
		cout << "Total registers per block: " << devProp.regsPerBlock << endl;
		cout << "Warp size: " << devProp.warpSize << endl;
		cout << "Maximum threads per block" << devProp.maxThreadsPerBlock << endl;

		for (int j = 0; j < 3; j++)
		{
			cout << "Maximum dimension " << j << " of block: " << devProp.maxThreadsDim[j] << endl;
			cout << "Maximum dimension " << j << " of grid: " << devProp.maxGridSize[j] << endl;
		}

		cout << "Clock rate: " << devProp.clockRate << endl;
		cout << "Total constant memory: " << devProp.totalConstMem << endl;
		cout << "Number of multiprocessors: " << devProp.multiProcessorCount << endl;
		cout << "\n" << endl;
	}

	cout << "Allocating memory on GPU..." << endl;
	cuda_alloc = cudaMalloc((void**) &device_universe,dev_size);
	cudaCheckAlloc(cuda_alloc,"Error allocating device universe");
	cuda_alloc = cudaMalloc((void**) &device_next_universe,dev_size);
	cudaCheckAlloc(cuda_alloc,"Error allocating device next universe");
	cuda_alloc = cudaMemcpy(device_universe,host_universe,dev_size,cudaMemcpyHostToDevice);
	cudaCheckAlloc(cuda_alloc,"Error copying universe to GPU");
	cuda_alloc = cudaMemcpy(device_next_universe,device_universe,dev_size,cudaMemcpyDeviceToDevice);
	cudaCheckAlloc(cuda_alloc,"Error copying next universe to GPU");
	cout << "Done!" << endl;

	cout << "Running game of life on GPU..." << endl;
	for (int i = 0; i < TICK_COUNT; i++)
	{
		tick <<< grid,threads >>>(device_universe,device_next_universe,rows,columns);
		cudaDeviceSynchronize();
		update <<< grid,threads >>>(device_universe,device_next_universe,rows,columns);
	}
	cout << "Done!" << endl;

	cout << "Storing result in main memory..." << endl;
	cudaMemcpy(host_universe,device_universe,dev_size,cudaMemcpyDeviceToHost);
	cout << "Done!" << endl;

	
	cudaFree(device_universe);
	cudaFree(device_next_universe);

	if (save_to_file)
	{
			cout << "Writing output file..." << endl;
			CreateDirectory(L"D:\\HPCTest",NULL);
			globalUniverse ->storeMatrix("D:\\HPCTest\\matrix.out");
			cout << "Done!" << endl;
	}
}